<!-- HTML header for doxygen 1.8.3.1-->
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml">
<head>
<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
<meta http-equiv="X-UA-Compatible" content="IE=9"/>
<meta name="generator" content="Doxygen 1.8.4"/>
<title>CUB: cub::WarpScan&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt; Class Template Reference</title>
<link href="tabs.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="jquery.js"></script>
<script type="text/javascript" src="dynsections.js"></script>
<link href="search/search.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="search/search.js"></script>
<script type="text/javascript">
  $(document).ready(function() { searchBox.OnSelectItem(0); });
</script>
<link href="doxygen.css" rel="stylesheet" type="text/css" />
<link href="extra_stylesheet.css" rel="stylesheet" type="text/css"/>
<link rel="shortcut icon" href="favicon.ico" type="image/x-icon" />
<script type="text/javascript">
  var _gaq = _gaq || [];
  _gaq.push(['_setAccount', 'UA-38890655-1']);
  _gaq.push(['_trackPageview']);
  (function() {
    var ga = document.createElement('script'); ga.type = 'text/javascript'; ga.async = true;
    ga.src = ('https:' == document.location.protocol ? 'https://ssl' : 'http://www') + '.google-analytics.com/ga.js';
    var s = document.getElementsByTagName('script')[0]; s.parentNode.insertBefore(ga, s);
  })();
</script>
</head>
<body>
<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
<div id="titlearea">
<table cellspacing="0" cellpadding="0">
 <tbody>
 <tr style="height: 56px;">
  <td style="padding-left: 0.5em;">
   <div id="projectname">CUB
   </div>
  </td>
 </tr>
 </tbody>
</table>
</div>
<!-- end header part -->
<!-- Generated by Doxygen 1.8.4 -->
<script type="text/javascript">
var searchBox = new SearchBox("searchBox", "search",false,'Search');
</script>
  <div id="navrow1" class="tabs">
    <ul class="tablist">
      <li><a href="index.html"><span>Main&#160;Page</span></a></li>
      <li><a href="modules.html"><span>Modules</span></a></li>
      <li class="current"><a href="annotated.html"><span>Classes</span></a></li>
      <li>
        <div id="MSearchBox" class="MSearchBoxInactive">
        <span class="left">
          <img id="MSearchSelect" src="search/mag_sel.png"
               onmouseover="return searchBox.OnSearchSelectShow()"
               onmouseout="return searchBox.OnSearchSelectHide()"
               alt=""/>
          <input type="text" id="MSearchField" value="Search" accesskey="S"
               onfocus="searchBox.OnSearchFieldFocus(true)" 
               onblur="searchBox.OnSearchFieldFocus(false)" 
               onkeyup="searchBox.OnSearchFieldChange(event)"/>
          </span><span class="right">
            <a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
          </span>
        </div>
      </li>
    </ul>
  </div>
  <div id="navrow2" class="tabs2">
    <ul class="tablist">
      <li><a href="annotated.html"><span>Class&#160;List</span></a></li>
      <li><a href="classes.html"><span>Class&#160;Index</span></a></li>
    </ul>
  </div>
<!-- window showing the filter options -->
<div id="MSearchSelectWindow"
     onmouseover="return searchBox.OnSearchSelectShow()"
     onmouseout="return searchBox.OnSearchSelectHide()"
     onkeydown="return searchBox.OnSearchSelectKey(event)">
<a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(0)"><span class="SelectionMark">&#160;</span>All</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(1)"><span class="SelectionMark">&#160;</span>Classes</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(2)"><span class="SelectionMark">&#160;</span>Namespaces</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(3)"><span class="SelectionMark">&#160;</span>Files</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(4)"><span class="SelectionMark">&#160;</span>Functions</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(5)"><span class="SelectionMark">&#160;</span>Variables</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(6)"><span class="SelectionMark">&#160;</span>Typedefs</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(7)"><span class="SelectionMark">&#160;</span>Enumerations</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(8)"><span class="SelectionMark">&#160;</span>Enumerator</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(9)"><span class="SelectionMark">&#160;</span>Groups</a></div>

<!-- iframe showing the search results (closed by default) -->
<div id="MSearchResultsWindow">
<iframe src="javascript:void(0)" frameborder="0" 
        name="MSearchResults" id="MSearchResults">
</iframe>
</div>

<div id="nav-path" class="navpath">
  <ul>
<li class="navelem"><a class="el" href="namespacecub.html">cub</a></li><li class="navelem"><a class="el" href="classcub_1_1_warp_scan.html">WarpScan</a></li>  </ul>
</div>
</div><!-- top -->
<div class="header">
  <div class="summary">
<a href="#nested-classes">Classes</a> &#124;
<a href="classcub_1_1_warp_scan-members.html">List of all members</a>  </div>
  <div class="headertitle">
<div class="title">cub::WarpScan&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt; Class Template Reference<div class="ingroups"><a class="el" href="group___warp_module.html">Warp-wide</a></div></div>  </div>
</div><!--header-->
<div class="contents">
<a name="details" id="details"></a><h2 class="groupheader">Detailed description</h2>
<div class="textblock"><h3>template&lt;
    typename T, 
    int LOGICAL_WARPS = 1, 
    int LOGICAL_WARP_THREADS = PtxArchProps::WARP_THREADS&gt;<br/>
class cub::WarpScan&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt;</h3>

<p>The <a class="el" href="classcub_1_1_warp_scan.html" title="The WarpScan class provides collective methods for computing a parallel prefix scan of items partitio...">WarpScan</a> class provides <a href="index.html#sec0"><em>collective</em></a> methods for computing a parallel prefix scan of items partitioned across CUDA warp threads. </p>
<div class="image">
<img src="warp_scan_logo.png" alt="warp_scan_logo.png"/>
<div class="caption">
.</div></div>
 <dl class="section user"><dt>Overview</dt><dd>Given a list of input elements and a binary reduction operator, a <a href="http://en.wikipedia.org/wiki/Prefix_sum"><em>prefix scan</em></a> produces an output list where each element is computed to be the reduction of the elements occurring earlier in the input list. <em>Prefix sum</em> connotes a prefix scan with the addition operator. The term <em>inclusive</em> indicates that the <em>i</em><sup>th</sup> output reduction incorporates the <em>i</em><sup>th</sup> input. The term <em>exclusive</em> indicates the <em>i</em><sup>th</sup> input is not incorporated into the <em>i</em><sup>th</sup> output reduction.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">T</td><td>The scan input/output element type </td></tr>
    <tr><td class="paramname">LOGICAL_WARPS</td><td><b>[optional]</b> The number of "logical" warps performing concurrent warp scans. Default is 1. </td></tr>
    <tr><td class="paramname">LOGICAL_WARP_THREADS</td><td><b>[optional]</b> The number of threads per "logical" warp (may be less than the number of hardware warp threads). Default is the warp size associated with the CUDA Compute Capability targeted by the compiler (e.g., 32 threads for SM20).</td></tr>
  </table>
  </dd>
</dl>
<dl class="section user"><dt>Simple Examples</dt><dd>Every thread in the warp uses the <a class="el" href="classcub_1_1_warp_scan.html" title="The WarpScan class provides collective methods for computing a parallel prefix scan of items partitio...">WarpScan</a> class by first specializing the <a class="el" href="classcub_1_1_warp_scan.html" title="The WarpScan class provides collective methods for computing a parallel prefix scan of items partitio...">WarpScan</a> type, then instantiating an instance with parameters for communication, and finally invoking collective member functions. </dd></dl>
<dl class="section user"><dt></dt><dd>The code snippet below illustrates four concurrent warp prefix sums within a block of 128 threads (one per each of the 32-thread warps). </dd></dl>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize WarpScan for 4 warps on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_warp_scan.html" title="The WarpScan class provides collective methods for computing a parallel prefix scan of items partitio...">cub::WarpScan&lt;int, 4&gt;</a> <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for WarpScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> WarpScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain one input item per thread</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data = ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Compute warp-wide prefix sums</span></div>
<div class="line">    <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>(temp_storage).ExclusiveSum(thread_data, thread_data);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>1, 1, 1, 1, ...</code>. The corresponding output <code>thread_data</code> in each of the four warps of threads will be <code>0, 1, 2, 3, ..., 31</code>.</dd></dl>
<dl class="section user"><dt></dt><dd>The code snippet below illustrates a single warp prefix sum within a block of 128 threads. </dd></dl>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize WarpScan for one warp on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_warp_scan.html" title="The WarpScan class provides collective methods for computing a parallel prefix scan of items partitio...">cub::WarpScan&lt;int, 1&gt;</a> <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for WarpScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> WarpScan::TempStorage temp_storage;</div>
<div class="line">    ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Only the first warp performs a prefix sum</span></div>
<div class="line">    <span class="keywordflow">if</span> (threadIdx.x &lt; 32)</div>
<div class="line">    {</div>
<div class="line">        <span class="comment">// Obtain one input item per thread</span></div>
<div class="line">        <span class="keywordtype">int</span> thread_data = ...</div>
<div class="line"></div>
<div class="line">        <span class="comment">// Compute warp-wide prefix sums</span></div>
<div class="line">        <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>(temp_storage).ExclusiveSum(thread_data, thread_data);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the warp of threads is <code>1, 1, 1, 1, ...</code>. The corresponding output <code>thread_data</code> will be <code>0, 1, 2, 3, ..., 31</code>.</dd></dl>
<dl class="section user"><dt>Usage and Performance Considerations</dt><dd><ul>
<li>Supports "logical" warps smaller than the physical warp size (e.g., a logical warp of 8 threads)</li>
<li>The number of entrant threads must be an multiple of <code>LOGICAL_WARP_THREADS</code> </li>
<li>Warp scans are concurrent if more than one warp is participating</li>
<li>Uses special instructions when applicable (e.g., warp <code>SHFL</code>)</li>
<li>Uses synchronization-free communication between warp lanes when applicable</li>
<li>Zero bank conflicts for most types.</li>
<li>Computation is slightly more efficient (i.e., having lower instruction overhead) for:<ul>
<li>Summation (<b><em>vs.</em></b> generic scan)</li>
<li>The architecture's warp size is a whole multiple of <code>LOGICAL_WARP_THREADS</code> </li>
</ul>
</li>
</ul>
</dd></dl>

<p>Definition at line <a class="el" href="warp__scan_8cuh_source.html#l00145">145</a> of file <a class="el" href="warp__scan_8cuh_source.html">warp_scan.cuh</a>.</p>
</div><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="nested-classes"></a>
Classes</h2></td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_warp_scan_1_1_temp_storage.html">TempStorage</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">The operations exposed by <a class="el" href="classcub_1_1_warp_scan.html" title="The WarpScan class provides collective methods for computing a parallel prefix scan of items partitio...">WarpScan</a> require a temporary memory allocation of this nested type for thread communication. This opaque storage can be allocated directly using the <code>__shared__</code> keyword. Alternatively, it can be aliased to externally allocated memory (shared or global) or <code>union</code>'d with other storage allocation types to facilitate memory reuse.  <a href="structcub_1_1_warp_scan_1_1_temp_storage.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="pub-methods"></a>
Public Methods</h2></td></tr>
<tr><td colspan="2"><div class="groupHeader">Collective constructors</div></td></tr>
<tr class="memitem:a27ca414b95f92f2b52701cda45273b05"><td class="memItemLeft" align="right" valign="top"><a class="anchor" id="a27ca414b95f92f2b52701cda45273b05"></a>
__device__ __forceinline__&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05">WarpScan</a> ()</td></tr>
<tr class="memdesc:a27ca414b95f92f2b52701cda45273b05"><td class="mdescLeft">&#160;</td><td class="mdescRight">Collective constructor for 1D thread blocks using a private static allocation of shared memory as temporary storage. Logical warp and lane identifiers are constructed from <code>threadIdx.x</code>. <br/></td></tr>
<tr class="separator:a27ca414b95f92f2b52701cda45273b05"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a037e31be8f0e394e2b7a6ce02f24042e"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_warp_scan.html#a037e31be8f0e394e2b7a6ce02f24042e">WarpScan</a> (<a class="el" href="structcub_1_1_warp_scan_1_1_temp_storage.html">TempStorage</a> &amp;temp_storage)</td></tr>
<tr class="memdesc:a037e31be8f0e394e2b7a6ce02f24042e"><td class="mdescLeft">&#160;</td><td class="mdescRight">Collective constructor for 1D thread blocks using the specified memory allocation as temporary storage. Logical warp and lane identifiers are constructed from <code>threadIdx.x</code>.  <a href="#a037e31be8f0e394e2b7a6ce02f24042e">More...</a><br/></td></tr>
<tr class="separator:a037e31be8f0e394e2b7a6ce02f24042e"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:aa8e1da90057f72813e2ad36af598dcc2"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_warp_scan.html#aa8e1da90057f72813e2ad36af598dcc2">WarpScan</a> (int warp_id, int lane_id)</td></tr>
<tr class="memdesc:aa8e1da90057f72813e2ad36af598dcc2"><td class="mdescLeft">&#160;</td><td class="mdescRight">Collective constructor using a private static allocation of shared memory as temporary storage. Threads are identified using the given warp and lane identifiers.  <a href="#aa8e1da90057f72813e2ad36af598dcc2">More...</a><br/></td></tr>
<tr class="separator:aa8e1da90057f72813e2ad36af598dcc2"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a9787551c8a6e2a23f9d32cd029151093"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_warp_scan.html#a9787551c8a6e2a23f9d32cd029151093">WarpScan</a> (<a class="el" href="structcub_1_1_warp_scan_1_1_temp_storage.html">TempStorage</a> &amp;temp_storage, int warp_id, int lane_id)</td></tr>
<tr class="memdesc:a9787551c8a6e2a23f9d32cd029151093"><td class="mdescLeft">&#160;</td><td class="mdescRight">Collective constructor using the specified memory allocation as temporary storage. Threads are identified using the given warp and lane identifiers.  <a href="#a9787551c8a6e2a23f9d32cd029151093">More...</a><br/></td></tr>
<tr class="separator:a9787551c8a6e2a23f9d32cd029151093"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr><td colspan="2"><div class="groupHeader">Inclusive prefix sums</div></td></tr>
<tr class="memitem:a64d246eb963f9a18d964e17a6f73fc27"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_warp_scan.html#a64d246eb963f9a18d964e17a6f73fc27">InclusiveSum</a> (T input, T &amp;output)</td></tr>
<tr class="memdesc:a64d246eb963f9a18d964e17a6f73fc27"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an inclusive prefix sum in each logical warp.  <a href="#a64d246eb963f9a18d964e17a6f73fc27">More...</a><br/></td></tr>
<tr class="separator:a64d246eb963f9a18d964e17a6f73fc27"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a601da756b6729894ec8a52080d920cd4"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_warp_scan.html#a601da756b6729894ec8a52080d920cd4">InclusiveSum</a> (T input, T &amp;output, T &amp;warp_aggregate)</td></tr>
<tr class="memdesc:a601da756b6729894ec8a52080d920cd4"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an inclusive prefix sum in each logical warp. Also provides every thread with the warp-wide <code>warp_aggregate</code> of all inputs.  <a href="#a601da756b6729894ec8a52080d920cd4">More...</a><br/></td></tr>
<tr class="separator:a601da756b6729894ec8a52080d920cd4"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a906d0d6381b3af1ad105ec69547350f3"><td class="memTemplParams" colspan="2">template&lt;typename WarpPrefixOp &gt; </td></tr>
<tr class="memitem:a906d0d6381b3af1ad105ec69547350f3"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_warp_scan.html#a906d0d6381b3af1ad105ec69547350f3">InclusiveSum</a> (T input, T &amp;output, T &amp;warp_aggregate, WarpPrefixOp &amp;warp_prefix_op)</td></tr>
<tr class="memdesc:a906d0d6381b3af1ad105ec69547350f3"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an inclusive prefix sum in each logical warp. Instead of using 0 as the warp-wide prefix, the call-back functor <code>warp_prefix_op</code> is invoked to provide the "seed" value that logically prefixes the warp's scan inputs. Also provides every thread with the warp-wide <code>warp_aggregate</code> of all inputs.  <a href="#a906d0d6381b3af1ad105ec69547350f3">More...</a><br/></td></tr>
<tr class="separator:a906d0d6381b3af1ad105ec69547350f3"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr><td colspan="2"><div class="groupHeader">Exclusive prefix sums</div></td></tr>
<tr class="memitem:acf4960fcb08e5960ca7c9aa28a34ae1c"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_warp_scan.html#acf4960fcb08e5960ca7c9aa28a34ae1c">ExclusiveSum</a> (T input, T &amp;output)</td></tr>
<tr class="memdesc:acf4960fcb08e5960ca7c9aa28a34ae1c"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an exclusive prefix sum in each logical warp.  <a href="#acf4960fcb08e5960ca7c9aa28a34ae1c">More...</a><br/></td></tr>
<tr class="separator:acf4960fcb08e5960ca7c9aa28a34ae1c"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ab68084364e2235b53998861731b267f8"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_warp_scan.html#ab68084364e2235b53998861731b267f8">ExclusiveSum</a> (T input, T &amp;output, T &amp;warp_aggregate)</td></tr>
<tr class="memdesc:ab68084364e2235b53998861731b267f8"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an exclusive prefix sum in each logical warp. Also provides every thread with the warp-wide <code>warp_aggregate</code> of all inputs.  <a href="#ab68084364e2235b53998861731b267f8">More...</a><br/></td></tr>
<tr class="separator:ab68084364e2235b53998861731b267f8"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a283236673ad2e54579f6b5ee97eac193"><td class="memTemplParams" colspan="2">template&lt;typename WarpPrefixOp &gt; </td></tr>
<tr class="memitem:a283236673ad2e54579f6b5ee97eac193"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_warp_scan.html#a283236673ad2e54579f6b5ee97eac193">ExclusiveSum</a> (T input, T &amp;output, T &amp;warp_aggregate, WarpPrefixOp &amp;warp_prefix_op)</td></tr>
<tr class="memdesc:a283236673ad2e54579f6b5ee97eac193"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an exclusive prefix sum in each logical warp. Instead of using 0 as the warp-wide prefix, the call-back functor <code>warp_prefix_op</code> is invoked to provide the "seed" value that logically prefixes the warp's scan inputs. Also provides every thread with the warp-wide <code>warp_aggregate</code> of all inputs.  <a href="#a283236673ad2e54579f6b5ee97eac193">More...</a><br/></td></tr>
<tr class="separator:a283236673ad2e54579f6b5ee97eac193"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr><td colspan="2"><div class="groupHeader">Inclusive prefix scans</div></td></tr>
<tr class="memitem:aafd14f3920c493b104a3a6a0cc6a39d2"><td class="memTemplParams" colspan="2">template&lt;typename ScanOp &gt; </td></tr>
<tr class="memitem:aafd14f3920c493b104a3a6a0cc6a39d2"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_warp_scan.html#aafd14f3920c493b104a3a6a0cc6a39d2">InclusiveScan</a> (T input, T &amp;output, ScanOp scan_op)</td></tr>
<tr class="memdesc:aafd14f3920c493b104a3a6a0cc6a39d2"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an inclusive prefix sum using the specified binary scan functor in each logical warp.  <a href="#aafd14f3920c493b104a3a6a0cc6a39d2">More...</a><br/></td></tr>
<tr class="separator:aafd14f3920c493b104a3a6a0cc6a39d2"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a6a2ac5f642393ccfc698aa515e192a22"><td class="memTemplParams" colspan="2">template&lt;typename ScanOp &gt; </td></tr>
<tr class="memitem:a6a2ac5f642393ccfc698aa515e192a22"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_warp_scan.html#a6a2ac5f642393ccfc698aa515e192a22">InclusiveScan</a> (T input, T &amp;output, ScanOp scan_op, T &amp;warp_aggregate)</td></tr>
<tr class="memdesc:a6a2ac5f642393ccfc698aa515e192a22"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an inclusive prefix sum using the specified binary scan functor in each logical warp. Also provides every thread with the warp-wide <code>warp_aggregate</code> of all inputs.  <a href="#a6a2ac5f642393ccfc698aa515e192a22">More...</a><br/></td></tr>
<tr class="separator:a6a2ac5f642393ccfc698aa515e192a22"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:adf74795dbe02947facd24657bc519c40"><td class="memTemplParams" colspan="2">template&lt;typename ScanOp , typename WarpPrefixOp &gt; </td></tr>
<tr class="memitem:adf74795dbe02947facd24657bc519c40"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_warp_scan.html#adf74795dbe02947facd24657bc519c40">InclusiveScan</a> (T input, T &amp;output, ScanOp scan_op, T &amp;warp_aggregate, WarpPrefixOp &amp;warp_prefix_op)</td></tr>
<tr class="memdesc:adf74795dbe02947facd24657bc519c40"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an inclusive prefix sum using the specified binary scan functor in each logical warp. The call-back functor <code>warp_prefix_op</code> is invoked to provide the "seed" value that logically prefixes the warp's scan inputs. Also provides every thread with the warp-wide <code>warp_aggregate</code> of all inputs.  <a href="#adf74795dbe02947facd24657bc519c40">More...</a><br/></td></tr>
<tr class="separator:adf74795dbe02947facd24657bc519c40"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr><td colspan="2"><div class="groupHeader">Exclusive prefix scans</div></td></tr>
<tr class="memitem:aafbc42fcea0a11b83289629554cbe625"><td class="memTemplParams" colspan="2">template&lt;typename ScanOp &gt; </td></tr>
<tr class="memitem:aafbc42fcea0a11b83289629554cbe625"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_warp_scan.html#aafbc42fcea0a11b83289629554cbe625">ExclusiveScan</a> (T input, T &amp;output, T identity, ScanOp scan_op)</td></tr>
<tr class="memdesc:aafbc42fcea0a11b83289629554cbe625"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an exclusive prefix scan using the specified binary scan functor in each logical warp.  <a href="#aafbc42fcea0a11b83289629554cbe625">More...</a><br/></td></tr>
<tr class="separator:aafbc42fcea0a11b83289629554cbe625"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a09d36c6cf146b69c1cf91d4781b96374"><td class="memTemplParams" colspan="2">template&lt;typename ScanOp &gt; </td></tr>
<tr class="memitem:a09d36c6cf146b69c1cf91d4781b96374"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_warp_scan.html#a09d36c6cf146b69c1cf91d4781b96374">ExclusiveScan</a> (T input, T &amp;output, T identity, ScanOp scan_op, T &amp;warp_aggregate)</td></tr>
<tr class="memdesc:a09d36c6cf146b69c1cf91d4781b96374"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an exclusive prefix scan using the specified binary scan functor in each logical warp. Also provides every thread with the warp-wide <code>warp_aggregate</code> of all inputs.  <a href="#a09d36c6cf146b69c1cf91d4781b96374">More...</a><br/></td></tr>
<tr class="separator:a09d36c6cf146b69c1cf91d4781b96374"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ada12695322d9e1f5354fcae9d0910113"><td class="memTemplParams" colspan="2">template&lt;typename ScanOp , typename WarpPrefixOp &gt; </td></tr>
<tr class="memitem:ada12695322d9e1f5354fcae9d0910113"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_warp_scan.html#ada12695322d9e1f5354fcae9d0910113">ExclusiveScan</a> (T input, T &amp;output, T identity, ScanOp scan_op, T &amp;warp_aggregate, WarpPrefixOp &amp;warp_prefix_op)</td></tr>
<tr class="memdesc:ada12695322d9e1f5354fcae9d0910113"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an exclusive prefix scan using the specified binary scan functor in each logical warp. The call-back functor <code>warp_prefix_op</code> is invoked to provide the "seed" value that logically prefixes the warp's scan inputs. Also provides every thread with the warp-wide <code>warp_aggregate</code> of all inputs.  <a href="#ada12695322d9e1f5354fcae9d0910113">More...</a><br/></td></tr>
<tr class="separator:ada12695322d9e1f5354fcae9d0910113"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr><td colspan="2"><div class="groupHeader">Identityless exclusive prefix scans</div></td></tr>
<tr class="memitem:a1f4e88a801bdfdfde342f23fb2768ea8"><td class="memTemplParams" colspan="2">template&lt;typename ScanOp &gt; </td></tr>
<tr class="memitem:a1f4e88a801bdfdfde342f23fb2768ea8"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_warp_scan.html#a1f4e88a801bdfdfde342f23fb2768ea8">ExclusiveScan</a> (T input, T &amp;output, ScanOp scan_op)</td></tr>
<tr class="memdesc:a1f4e88a801bdfdfde342f23fb2768ea8"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an exclusive prefix scan using the specified binary scan functor in each logical warp. Because no identity value is supplied, the <code>output</code> computed for <em>warp-lane</em><sub>0</sub> is undefined.  <a href="#a1f4e88a801bdfdfde342f23fb2768ea8">More...</a><br/></td></tr>
<tr class="separator:a1f4e88a801bdfdfde342f23fb2768ea8"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:af51aef1f68a2ca876f7515b6b3213d72"><td class="memTemplParams" colspan="2">template&lt;typename ScanOp &gt; </td></tr>
<tr class="memitem:af51aef1f68a2ca876f7515b6b3213d72"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_warp_scan.html#af51aef1f68a2ca876f7515b6b3213d72">ExclusiveScan</a> (T input, T &amp;output, ScanOp scan_op, T &amp;warp_aggregate)</td></tr>
<tr class="memdesc:af51aef1f68a2ca876f7515b6b3213d72"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an exclusive prefix scan using the specified binary scan functor in each logical warp. Because no identity value is supplied, the <code>output</code> computed for <em>warp-lane</em><sub>0</sub> is undefined. Also provides every thread with the warp-wide <code>warp_aggregate</code> of all inputs.  <a href="#af51aef1f68a2ca876f7515b6b3213d72">More...</a><br/></td></tr>
<tr class="separator:af51aef1f68a2ca876f7515b6b3213d72"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a024d74c48e45b489843053097b0638d1"><td class="memTemplParams" colspan="2">template&lt;typename ScanOp , typename WarpPrefixOp &gt; </td></tr>
<tr class="memitem:a024d74c48e45b489843053097b0638d1"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_warp_scan.html#a024d74c48e45b489843053097b0638d1">ExclusiveScan</a> (T input, T &amp;output, ScanOp scan_op, T &amp;warp_aggregate, WarpPrefixOp &amp;warp_prefix_op)</td></tr>
<tr class="memdesc:a024d74c48e45b489843053097b0638d1"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an exclusive prefix scan using the specified binary scan functor in each logical warp. The <code>warp_prefix_op</code> value from thread-thread-lane<sub>0</sub> is applied to all scan outputs. Also computes the warp-wide <code>warp_aggregate</code> of all inputs for thread-thread-lane<sub>0</sub>.  <a href="#a024d74c48e45b489843053097b0638d1">More...</a><br/></td></tr>
<tr class="separator:a024d74c48e45b489843053097b0638d1"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table>
<h2 class="groupheader">Constructor &amp; Destructor Documentation</h2>
<a class="anchor" id="a037e31be8f0e394e2b7a6ce02f24042e"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int LOGICAL_WARPS = 1, int LOGICAL_WARP_THREADS = PtxArchProps::WARP_THREADS&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ <a class="el" href="classcub_1_1_warp_scan.html">cub::WarpScan</a>&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt;::<a class="el" href="classcub_1_1_warp_scan.html">WarpScan</a> </td>
          <td>(</td>
          <td class="paramtype"><a class="el" href="structcub_1_1_warp_scan_1_1_temp_storage.html">TempStorage</a> &amp;&#160;</td>
          <td class="paramname"><em>temp_storage</em>)</td><td></td>
          <td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Collective constructor for 1D thread blocks using the specified memory allocation as temporary storage. Logical warp and lane identifiers are constructed from <code>threadIdx.x</code>. </p>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">temp_storage</td><td>Reference to memory allocation having layout type <a class="el" href="structcub_1_1_warp_scan_1_1_temp_storage.html" title="The operations exposed by WarpScan require a temporary memory allocation of this nested type for thre...">TempStorage</a> </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="warp__scan_8cuh_source.html#l00222">222</a> of file <a class="el" href="warp__scan_8cuh_source.html">warp_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="aa8e1da90057f72813e2ad36af598dcc2"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int LOGICAL_WARPS = 1, int LOGICAL_WARP_THREADS = PtxArchProps::WARP_THREADS&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ <a class="el" href="classcub_1_1_warp_scan.html">cub::WarpScan</a>&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt;::<a class="el" href="classcub_1_1_warp_scan.html">WarpScan</a> </td>
          <td>(</td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>warp_id</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>lane_id</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Collective constructor using a private static allocation of shared memory as temporary storage. Threads are identified using the given warp and lane identifiers. </p>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">warp_id</td><td>A suitable warp membership identifier </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">lane_id</td><td>A lane identifier within the warp </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="warp__scan_8cuh_source.html#l00238">238</a> of file <a class="el" href="warp__scan_8cuh_source.html">warp_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a9787551c8a6e2a23f9d32cd029151093"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int LOGICAL_WARPS = 1, int LOGICAL_WARP_THREADS = PtxArchProps::WARP_THREADS&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ <a class="el" href="classcub_1_1_warp_scan.html">cub::WarpScan</a>&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt;::<a class="el" href="classcub_1_1_warp_scan.html">WarpScan</a> </td>
          <td>(</td>
          <td class="paramtype"><a class="el" href="structcub_1_1_warp_scan_1_1_temp_storage.html">TempStorage</a> &amp;&#160;</td>
          <td class="paramname"><em>temp_storage</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>warp_id</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>lane_id</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Collective constructor using the specified memory allocation as temporary storage. Threads are identified using the given warp and lane identifiers. </p>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">temp_storage</td><td>Reference to memory allocation having layout type <a class="el" href="structcub_1_1_warp_scan_1_1_temp_storage.html" title="The operations exposed by WarpScan require a temporary memory allocation of this nested type for thre...">TempStorage</a> </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">warp_id</td><td>A suitable warp membership identifier </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">lane_id</td><td>A lane identifier within the warp </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="warp__scan_8cuh_source.html#l00251">251</a> of file <a class="el" href="warp__scan_8cuh_source.html">warp_scan.cuh</a>.</p>

</div>
</div>
<h2 class="groupheader">Member Function Documentation</h2>
<a class="anchor" id="aafbc42fcea0a11b83289629554cbe625"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int LOGICAL_WARPS = 1, int LOGICAL_WARP_THREADS = PtxArchProps::WARP_THREADS&gt; </div>
<div class="memtemplate">
template&lt;typename ScanOp &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_warp_scan.html">cub::WarpScan</a>&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt;::ExclusiveScan </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>output</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>identity</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">ScanOp&#160;</td>
          <td class="paramname"><em>scan_op</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an exclusive prefix scan using the specified binary scan functor in each logical warp. </p>
<p>Supports non-commutative scan operators.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates four concurrent warp-wide exclusive prefix max scans within a block of 128 threads (one per each of the 32-thread warps). </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize WarpScan for 4 warps on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_warp_scan.html" title="The WarpScan class provides collective methods for computing a parallel prefix scan of items partitio...">cub::WarpScan&lt;int, 4&gt;</a> <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for WarpScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> WarpScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain one input item per thread</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data = ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Compute exclusive warp-wide prefix max scans</span></div>
<div class="line">    <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>(temp_storage).ExclusiveScan(thread_data, thread_data, INT_MIN, <a class="code" href="structcub_1_1_max.html" title="Default max functor. ">cub::Max</a>());</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>0, -1, 2, -3, ..., 126, -127</code>. The corresponding output <code>thread_data</code> in the first warp would be <code>INT_MIN, 0, 0, 2, ..., 28, 30</code>, the output for the second warp would be <code>30, 32, 32, 34, ..., 60, 62</code>, etc.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ScanOp</td><td><b>[inferred]</b> Binary scan operator type having member <code>T operator()(const T &amp;a, const T &amp;b)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input item. </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output item. May be aliased with <code>input</code>. </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">identity</td><td>Identity value </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">scan_op</td><td>Binary scan operator </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="warp__scan_8cuh_source.html#l00931">931</a> of file <a class="el" href="warp__scan_8cuh_source.html">warp_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a09d36c6cf146b69c1cf91d4781b96374"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int LOGICAL_WARPS = 1, int LOGICAL_WARP_THREADS = PtxArchProps::WARP_THREADS&gt; </div>
<div class="memtemplate">
template&lt;typename ScanOp &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_warp_scan.html">cub::WarpScan</a>&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt;::ExclusiveScan </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>output</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>identity</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">ScanOp&#160;</td>
          <td class="paramname"><em>scan_op</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>warp_aggregate</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an exclusive prefix scan using the specified binary scan functor in each logical warp. Also provides every thread with the warp-wide <code>warp_aggregate</code> of all inputs. </p>
<p>Supports non-commutative scan operators.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates four concurrent warp-wide exclusive prefix max scans within a block of 128 threads (one per each of the 32-thread warps). </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize WarpScan for 4 warps on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_warp_scan.html" title="The WarpScan class provides collective methods for computing a parallel prefix scan of items partitio...">cub::WarpScan&lt;int, 4&gt;</a> <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for WarpScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> WarpScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain one input item per thread</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data = ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Compute exclusive warp-wide prefix max scans</span></div>
<div class="line">    <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>(temp_storage).ExclusiveScan(thread_data, thread_data, INT_MIN, <a class="code" href="structcub_1_1_max.html" title="Default max functor. ">cub::Max</a>());</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>0, -1, 2, -3, ..., 126, -127</code>. The corresponding output <code>thread_data</code> in the first warp would be <code>INT_MIN, 0, 0, 2, ..., 28, 30</code>, the output for the second warp would be <code>30, 32, 32, 34, ..., 60, 62</code>, etc. Furthermore, <code>warp_aggregate</code> would be assigned <code>30</code> for threads in the first warp, <code>62</code> for threads in the second warp, etc.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ScanOp</td><td><b>[inferred]</b> Binary scan operator type having member <code>T operator()(const T &amp;a, const T &amp;b)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input item. </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output item. May be aliased with <code>input</code>. </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">identity</td><td>Identity value </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">scan_op</td><td>Binary scan operator </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">warp_aggregate</td><td>Warp-wide aggregate reduction of input items. </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="warp__scan_8cuh_source.html#l00979">979</a> of file <a class="el" href="warp__scan_8cuh_source.html">warp_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="ada12695322d9e1f5354fcae9d0910113"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int LOGICAL_WARPS = 1, int LOGICAL_WARP_THREADS = PtxArchProps::WARP_THREADS&gt; </div>
<div class="memtemplate">
template&lt;typename ScanOp , typename WarpPrefixOp &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_warp_scan.html">cub::WarpScan</a>&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt;::ExclusiveScan </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>output</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>identity</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">ScanOp&#160;</td>
          <td class="paramname"><em>scan_op</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>warp_aggregate</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">WarpPrefixOp &amp;&#160;</td>
          <td class="paramname"><em>warp_prefix_op</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an exclusive prefix scan using the specified binary scan functor in each logical warp. The call-back functor <code>warp_prefix_op</code> is invoked to provide the "seed" value that logically prefixes the warp's scan inputs. Also provides every thread with the warp-wide <code>warp_aggregate</code> of all inputs. </p>
<p>The <code>warp_prefix_op</code> functor must implement a member function <code>T operator()(T warp_aggregate)</code>. The functor's input parameter <code>warp_aggregate</code> is the same value also returned by the scan operation. The functor will be invoked by the entire warp of threads, however only the return value from <em>lane</em><sub>0</sub> is applied as the threadblock-wide prefix. Can be stateful.</p>
<p>Supports non-commutative scan operators.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates a single thread block of 32 threads (one warp) that progressively computes an exclusive prefix max scan over multiple "tiles" of input using a prefix functor to maintain a running total between block-wide scans. Each tile consists of 32 integer items that are partitioned across the warp. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line"><span class="comment">// A stateful callback functor that maintains a running prefix to be applied</span></div>
<div class="line"><span class="comment">// during consecutive scan operations.</span></div>
<div class="line"><span class="keyword">struct </span>WarpPrefixOp</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Running prefix</span></div>
<div class="line">    <span class="keywordtype">int</span> running_total;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Constructor</span></div>
<div class="line">    __device__ WarpPrefixOp(<span class="keywordtype">int</span> running_total) : running_total(running_total) {}</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Callback operator to be entered by the entire warp. Lane-0 is responsible</span></div>
<div class="line">    <span class="comment">// for returning a value for seeding the warp-wide scan.</span></div>
<div class="line">    __device__ <span class="keywordtype">int</span> operator()(<span class="keywordtype">int</span> warp_aggregate)</div>
<div class="line">    {</div>
<div class="line">        <span class="keywordtype">int</span> old_prefix = running_total;</div>
<div class="line">        running_total = (warp_aggregate &gt; old_prefix) ? warp_aggregate : old_prefix;</div>
<div class="line">        <span class="keywordflow">return</span> old_prefix;</div>
<div class="line">    }</div>
<div class="line">};</div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(<span class="keywordtype">int</span> *d_data, <span class="keywordtype">int</span> num_items, ...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize WarpScan for one warp</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_warp_scan.html" title="The WarpScan class provides collective methods for computing a parallel prefix scan of items partitio...">cub::WarpScan&lt;int, 1&gt;</a> <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for WarpScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> WarpScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Initialize running total</span></div>
<div class="line">    WarpPrefixOp prefix_op(INT_MIN);</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Have the warp iterate over segments of items</span></div>
<div class="line">    <span class="keywordflow">for</span> (<span class="keywordtype">int</span> block_offset = 0; block_offset &lt; num_items; block_offset += 32)</div>
<div class="line">    {</div>
<div class="line">        <span class="comment">// Load a segment of consecutive items</span></div>
<div class="line">        <span class="keywordtype">int</span> thread_data = d_data[block_offset];</div>
<div class="line"></div>
<div class="line">        <span class="comment">// Collectively compute the warp-wide exclusive prefix max scan</span></div>
<div class="line">        <span class="keywordtype">int</span> warp_aggregate;</div>
<div class="line">        <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>(temp_storage).ExclusiveScan(</div>
<div class="line">            thread_data, thread_data, INT_MIN, <a class="code" href="structcub_1_1_max.html" title="Default max functor. ">cub::Max</a>(), warp_aggregate, prefix_op);</div>
<div class="line"></div>
<div class="line">        <span class="comment">// Store scanned items to output segment</span></div>
<div class="line">        d_data[block_offset] = thread_data;</div>
<div class="line">    }</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the input <code>d_data</code> is <code>0, -1, 2, -3, 4, -5, ...</code>. The corresponding output for the first segment will be <code>INT_MIN, 0, 0, 2, ..., 28, 30</code>. The output for the second segment will be <code>30, 32, 32, 34, ..., 60, 62</code>. Furthermore, <code>block_aggregate</code> will be assigned <code>30</code> in all threads after the first scan, assigned <code>62</code> after the second scan, etc.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ScanOp</td><td><b>[inferred]</b> Binary scan operator type having member <code>T operator()(const T &amp;a, const T &amp;b)</code> </td></tr>
    <tr><td class="paramname">WarpPrefixOp</td><td><b>[inferred]</b> Call-back functor type having member <code>T operator()(T warp_aggregate)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input item. </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output item. May be aliased with <code>input</code>. </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">identity</td><td>Identity value </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">scan_op</td><td>Binary scan operator </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">warp_aggregate</td><td><b>[<em>warp-lane</em><sub>0</sub> only]</b> Warp-wide aggregate reduction of input items (exclusive of the <code>warp_prefix_op</code> value). </td></tr>
    <tr><td class="paramdir">[in,out]</td><td class="paramname">warp_prefix_op</td><td><b>[<em>warp-lane</em><sub>0</sub> only]</b> Call-back functor for specifying a warp-wide prefix to be applied to all inputs. </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="warp__scan_8cuh_source.html#l01069">1069</a> of file <a class="el" href="warp__scan_8cuh_source.html">warp_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a1f4e88a801bdfdfde342f23fb2768ea8"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int LOGICAL_WARPS = 1, int LOGICAL_WARP_THREADS = PtxArchProps::WARP_THREADS&gt; </div>
<div class="memtemplate">
template&lt;typename ScanOp &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_warp_scan.html">cub::WarpScan</a>&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt;::ExclusiveScan </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>output</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">ScanOp&#160;</td>
          <td class="paramname"><em>scan_op</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an exclusive prefix scan using the specified binary scan functor in each logical warp. Because no identity value is supplied, the <code>output</code> computed for <em>warp-lane</em><sub>0</sub> is undefined. </p>
<p>Supports non-commutative scan operators.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates four concurrent warp-wide exclusive prefix max scans within a block of 128 threads (one per each of the 32-thread warps). </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize WarpScan for 4 warps on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_warp_scan.html" title="The WarpScan class provides collective methods for computing a parallel prefix scan of items partitio...">cub::WarpScan&lt;int, 4&gt;</a> <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for WarpScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> WarpScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain one input item per thread</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data = ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Compute exclusive warp-wide prefix max scans</span></div>
<div class="line">    <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>(temp_storage).ExclusiveScan(thread_data, thread_data, <a class="code" href="structcub_1_1_max.html" title="Default max functor. ">cub::Max</a>());</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>0, -1, 2, -3, ..., 126, -127</code>. The corresponding output <code>thread_data</code> in the first warp would be <code>?, 0, 0, 2, ..., 28, 30</code>, the output for the second warp would be <code>?, 32, 32, 34, ..., 60, 62</code>, etc. (The output <code>thread_data</code> in each warp lane0 is undefined.)</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ScanOp</td><td><b>[inferred]</b> Binary scan operator type having member <code>T operator()(const T &amp;a, const T &amp;b)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input item. </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output item. May be aliased with <code>input</code>. </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">scan_op</td><td>Binary scan operator </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="warp__scan_8cuh_source.html#l01135">1135</a> of file <a class="el" href="warp__scan_8cuh_source.html">warp_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="af51aef1f68a2ca876f7515b6b3213d72"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int LOGICAL_WARPS = 1, int LOGICAL_WARP_THREADS = PtxArchProps::WARP_THREADS&gt; </div>
<div class="memtemplate">
template&lt;typename ScanOp &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_warp_scan.html">cub::WarpScan</a>&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt;::ExclusiveScan </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>output</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">ScanOp&#160;</td>
          <td class="paramname"><em>scan_op</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>warp_aggregate</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an exclusive prefix scan using the specified binary scan functor in each logical warp. Because no identity value is supplied, the <code>output</code> computed for <em>warp-lane</em><sub>0</sub> is undefined. Also provides every thread with the warp-wide <code>warp_aggregate</code> of all inputs. </p>
<p>Supports non-commutative scan operators.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates four concurrent warp-wide exclusive prefix max scans within a block of 128 threads (one per each of the 32-thread warps). </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize WarpScan for 4 warps on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_warp_scan.html" title="The WarpScan class provides collective methods for computing a parallel prefix scan of items partitio...">cub::WarpScan&lt;int, 4&gt;</a> <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for WarpScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> WarpScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain one input item per thread</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data = ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Compute exclusive warp-wide prefix max scans</span></div>
<div class="line">    <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>(temp_storage).ExclusiveScan(thread_data, thread_data, <a class="code" href="structcub_1_1_max.html" title="Default max functor. ">cub::Max</a>());</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>0, -1, 2, -3, ..., 126, -127</code>. The corresponding output <code>thread_data</code> in the first warp would be <code>?, 0, 0, 2, ..., 28, 30</code>, the output for the second warp would be <code>?, 32, 32, 34, ..., 60, 62</code>, etc. (The output <code>thread_data</code> in each warp lane0 is undefined.) Furthermore, <code>warp_aggregate</code> would be assigned <code>30</code> for threads in the first warp, <code>62</code> for threads in the second warp, etc.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ScanOp</td><td><b>[inferred]</b> Binary scan operator type having member <code>T operator()(const T &amp;a, const T &amp;b)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input item. </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output item. May be aliased with <code>input</code>. </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">scan_op</td><td>Binary scan operator </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">warp_aggregate</td><td>Warp-wide aggregate reduction of input items. </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="warp__scan_8cuh_source.html#l01182">1182</a> of file <a class="el" href="warp__scan_8cuh_source.html">warp_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a024d74c48e45b489843053097b0638d1"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int LOGICAL_WARPS = 1, int LOGICAL_WARP_THREADS = PtxArchProps::WARP_THREADS&gt; </div>
<div class="memtemplate">
template&lt;typename ScanOp , typename WarpPrefixOp &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_warp_scan.html">cub::WarpScan</a>&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt;::ExclusiveScan </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>output</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">ScanOp&#160;</td>
          <td class="paramname"><em>scan_op</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>warp_aggregate</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">WarpPrefixOp &amp;&#160;</td>
          <td class="paramname"><em>warp_prefix_op</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an exclusive prefix scan using the specified binary scan functor in each logical warp. The <code>warp_prefix_op</code> value from thread-thread-lane<sub>0</sub> is applied to all scan outputs. Also computes the warp-wide <code>warp_aggregate</code> of all inputs for thread-thread-lane<sub>0</sub>. </p>
<p>The <code>warp_prefix_op</code> functor must implement a member function <code>T operator()(T warp_aggregate)</code>. The functor's input parameter <code>warp_aggregate</code> is the same value also returned by the scan operation. The functor will be invoked by the entire warp of threads, however only the return value from <em>lane</em><sub>0</sub> is applied as the threadblock-wide prefix. Can be stateful.</p>
<p>Supports non-commutative scan operators.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates a single thread block of 32 threads (one warp) that progressively computes an exclusive prefix max scan over multiple "tiles" of input using a prefix functor to maintain a running total between block-wide scans. Each tile consists of 32 integer items that are partitioned across the warp. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line"><span class="comment">// A stateful callback functor that maintains a running prefix to be applied</span></div>
<div class="line"><span class="comment">// during consecutive scan operations.</span></div>
<div class="line"><span class="keyword">struct </span>WarpPrefixOp</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Running prefix</span></div>
<div class="line">    <span class="keywordtype">int</span> running_total;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Constructor</span></div>
<div class="line">    __device__ WarpPrefixOp(<span class="keywordtype">int</span> running_total) : running_total(running_total) {}</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Callback operator to be entered by the entire warp. Lane-0 is responsible</span></div>
<div class="line">    <span class="comment">// for returning a value for seeding the warp-wide scan.</span></div>
<div class="line">    __device__ <span class="keywordtype">int</span> operator()(<span class="keywordtype">int</span> warp_aggregate)</div>
<div class="line">    {</div>
<div class="line">        <span class="keywordtype">int</span> old_prefix = running_total;</div>
<div class="line">        running_total = (warp_aggregate &gt; old_prefix) ? warp_aggregate : old_prefix;</div>
<div class="line">        <span class="keywordflow">return</span> old_prefix;</div>
<div class="line">    }</div>
<div class="line">};</div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(<span class="keywordtype">int</span> *d_data, <span class="keywordtype">int</span> num_items, ...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize WarpScan for one warp</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_warp_scan.html" title="The WarpScan class provides collective methods for computing a parallel prefix scan of items partitio...">cub::WarpScan&lt;int, 1&gt;</a> <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for WarpScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> WarpScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Initialize running total</span></div>
<div class="line">    WarpPrefixOp prefix_op(INT_MIN);</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Have the warp iterate over segments of items</span></div>
<div class="line">    <span class="keywordflow">for</span> (<span class="keywordtype">int</span> block_offset = 0; block_offset &lt; num_items; block_offset += 32)</div>
<div class="line">    {</div>
<div class="line">        <span class="comment">// Load a segment of consecutive items</span></div>
<div class="line">        <span class="keywordtype">int</span> thread_data = d_data[block_offset];</div>
<div class="line"></div>
<div class="line">        <span class="comment">// Collectively compute the warp-wide exclusive prefix max scan</span></div>
<div class="line">        <span class="keywordtype">int</span> warp_aggregate;</div>
<div class="line">        <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>(temp_storage).ExclusiveScan(</div>
<div class="line">            thread_data, thread_data, INT_MIN, <a class="code" href="structcub_1_1_max.html" title="Default max functor. ">cub::Max</a>(), warp_aggregate, prefix_op);</div>
<div class="line"></div>
<div class="line">        <span class="comment">// Store scanned items to output segment</span></div>
<div class="line">        d_data[block_offset] = thread_data;</div>
<div class="line">    }</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the input <code>d_data</code> is <code>0, -1, 2, -3, 4, -5, ...</code>. The corresponding output for the first segment will be <code>INT_MIN, 0, 0, 2, ..., 28, 30</code>. The output for the second segment will be <code>30, 32, 32, 34, ..., 60, 62</code>. Furthermore, <code>block_aggregate</code> will be assigned <code>30</code> in all threads after the first scan, assigned <code>62</code> after the second scan, etc.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ScanOp</td><td><b>[inferred]</b> Binary scan operator type having member <code>T operator()(const T &amp;a, const T &amp;b)</code> </td></tr>
    <tr><td class="paramname">WarpPrefixOp</td><td><b>[inferred]</b> Call-back functor type having member <code>T operator()(T warp_aggregate)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input item. </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output item. May be aliased with <code>input</code>. </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">scan_op</td><td>Binary scan operator </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">warp_aggregate</td><td><b>[<em>warp-lane</em><sub>0</sub> only]</b> Warp-wide aggregate reduction of input items (exclusive of the <code>warp_prefix_op</code> value). </td></tr>
    <tr><td class="paramdir">[in,out]</td><td class="paramname">warp_prefix_op</td><td><b>[<em>warp-lane</em><sub>0</sub> only]</b> Call-back functor for specifying a warp-wide prefix to be applied to all inputs. </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="warp__scan_8cuh_source.html#l01271">1271</a> of file <a class="el" href="warp__scan_8cuh_source.html">warp_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="acf4960fcb08e5960ca7c9aa28a34ae1c"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int LOGICAL_WARPS = 1, int LOGICAL_WARP_THREADS = PtxArchProps::WARP_THREADS&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_warp_scan.html">cub::WarpScan</a>&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt;::ExclusiveSum </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>output</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an exclusive prefix sum in each logical warp. </p>
<p>This operation assumes the value of obtained by the <code>T</code>'s default constructor (or by zero-initialization if no user-defined default constructor exists) is suitable as the identity value "zero" for addition.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates four concurrent warp-wide exclusive prefix sums within a block of 128 threads (one per each of the 32-thread warps). </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize WarpScan for 4 warps on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_warp_scan.html" title="The WarpScan class provides collective methods for computing a parallel prefix scan of items partitio...">cub::WarpScan&lt;int, 4&gt;</a> <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for WarpScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> WarpScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain one input item per thread</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data = ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Compute exclusive warp-wide prefix sums</span></div>
<div class="line">    <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>(temp_storage).ExclusiveSum(thread_data, thread_data);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>1, 1, 1, 1, ...</code>. The corresponding output <code>thread_data</code> in each of the four warps of threads will be <code>0, 1, 2, ..., 31</code>. </dd></dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input item. </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output item. May be aliased with <code>input</code>. </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="warp__scan_8cuh_source.html#l00547">547</a> of file <a class="el" href="warp__scan_8cuh_source.html">warp_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="ab68084364e2235b53998861731b267f8"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int LOGICAL_WARPS = 1, int LOGICAL_WARP_THREADS = PtxArchProps::WARP_THREADS&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_warp_scan.html">cub::WarpScan</a>&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt;::ExclusiveSum </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>output</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>warp_aggregate</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an exclusive prefix sum in each logical warp. Also provides every thread with the warp-wide <code>warp_aggregate</code> of all inputs. </p>
<p>This operation assumes the value of obtained by the <code>T</code>'s default constructor (or by zero-initialization if no user-defined default constructor exists) is suitable as the identity value "zero" for addition.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates four concurrent warp-wide exclusive prefix sums within a block of 128 threads (one per each of the 32-thread warps). </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize WarpScan for 4 warps on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_warp_scan.html" title="The WarpScan class provides collective methods for computing a parallel prefix scan of items partitio...">cub::WarpScan&lt;int, 4&gt;</a> <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for WarpScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> WarpScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain one input item per thread</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data = ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Compute exclusive warp-wide prefix sums</span></div>
<div class="line">    <span class="keywordtype">int</span> warp_aggregate;</div>
<div class="line">    <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>(temp_storage).ExclusiveSum(thread_data, thread_data, warp_aggregate);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>1, 1, 1, 1, ...</code>. The corresponding output <code>thread_data</code> in each of the four warps of threads will be <code>0, 1, 2, ..., 31</code>. Furthermore, <code>warp_aggregate</code> for all threads in all warps will be <code>32</code>. </dd></dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input item. </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output item. May be aliased with <code>input</code>. </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">warp_aggregate</td><td>Warp-wide aggregate reduction of input items. </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="warp__scan_8cuh_source.html#l00592">592</a> of file <a class="el" href="warp__scan_8cuh_source.html">warp_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a283236673ad2e54579f6b5ee97eac193"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int LOGICAL_WARPS = 1, int LOGICAL_WARP_THREADS = PtxArchProps::WARP_THREADS&gt; </div>
<div class="memtemplate">
template&lt;typename WarpPrefixOp &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_warp_scan.html">cub::WarpScan</a>&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt;::ExclusiveSum </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>output</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>warp_aggregate</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">WarpPrefixOp &amp;&#160;</td>
          <td class="paramname"><em>warp_prefix_op</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an exclusive prefix sum in each logical warp. Instead of using 0 as the warp-wide prefix, the call-back functor <code>warp_prefix_op</code> is invoked to provide the "seed" value that logically prefixes the warp's scan inputs. Also provides every thread with the warp-wide <code>warp_aggregate</code> of all inputs. </p>
<p>This operation assumes the value of obtained by the <code>T</code>'s default constructor (or by zero-initialization if no user-defined default constructor exists) is suitable as the identity value "zero" for addition.</p>
<p>The <code>warp_prefix_op</code> functor must implement a member function <code>T operator()(T warp_aggregate)</code>. The functor's input parameter <code>warp_aggregate</code> is the same value also returned by the scan operation. The functor will be invoked by the entire warp of threads, however only the return value from <em>lane</em><sub>0</sub> is applied as the threadblock-wide prefix. Can be stateful.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates a single thread block of 32 threads (one warp) that progressively computes an exclusive prefix sum over multiple "tiles" of input using a prefix functor to maintain a running total between block-wide scans. Each tile consists of 32 integer items that are partitioned across the warp. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line"><span class="comment">// A stateful callback functor that maintains a running prefix to be applied</span></div>
<div class="line"><span class="comment">// during consecutive scan operations.</span></div>
<div class="line"><span class="keyword">struct </span>WarpPrefixOp</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Running prefix</span></div>
<div class="line">    <span class="keywordtype">int</span> running_total;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Constructor</span></div>
<div class="line">    __device__ WarpPrefixOp(<span class="keywordtype">int</span> running_total) : running_total(running_total) {}</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Callback operator to be entered by the entire warp. Lane-0 is responsible</span></div>
<div class="line">    <span class="comment">// for returning a value for seeding the warp-wide scan.</span></div>
<div class="line">    __device__ <span class="keywordtype">int</span> operator()(<span class="keywordtype">int</span> warp_aggregate)</div>
<div class="line">    {</div>
<div class="line">        <span class="keywordtype">int</span> old_prefix = running_total;</div>
<div class="line">        running_total += warp_aggregate;</div>
<div class="line">        <span class="keywordflow">return</span> old_prefix;</div>
<div class="line">    }</div>
<div class="line">};</div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(<span class="keywordtype">int</span> *d_data, <span class="keywordtype">int</span> num_items, ...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize WarpScan for one warp</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_warp_scan.html" title="The WarpScan class provides collective methods for computing a parallel prefix scan of items partitio...">cub::WarpScan&lt;int, 1&gt;</a> <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for WarpScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> WarpScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Initialize running total</span></div>
<div class="line">    WarpPrefixOp prefix_op(0);</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Have the warp iterate over segments of items</span></div>
<div class="line">    <span class="keywordflow">for</span> (<span class="keywordtype">int</span> block_offset = 0; block_offset &lt; num_items; block_offset += 32)</div>
<div class="line">    {</div>
<div class="line">        <span class="comment">// Load a segment of consecutive items</span></div>
<div class="line">        <span class="keywordtype">int</span> thread_data = d_data[block_offset];</div>
<div class="line"></div>
<div class="line">        <span class="comment">// Collectively compute the warp-wide exclusive prefix sum</span></div>
<div class="line">        <span class="keywordtype">int</span> warp_aggregate;</div>
<div class="line">        <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>(temp_storage).ExclusiveSum(</div>
<div class="line">            thread_data, thread_data, warp_aggregate, prefix_op);</div>
<div class="line"></div>
<div class="line">        <span class="comment">// Store scanned items to output segment</span></div>
<div class="line">        d_data[block_offset] = thread_data;</div>
<div class="line">    }</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the input <code>d_data</code> is <code>1, 1, 1, 1, 1, 1, 1, 1, ...</code>. The corresponding output for the first segment will be <code>0, 1, 2, ..., 31</code>. The output for the second segment will be <code>32, 33, 34, ..., 63</code>. Furthermore, the value <code>32</code> will be stored in <code>warp_aggregate</code> for all threads after each scan.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">WarpPrefixOp</td><td><b>[inferred]</b> Call-back functor type having member <code>T operator()(T warp_aggregate)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input item. </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output item. May be aliased with <code>input</code>. </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">warp_aggregate</td><td><b>[<em>warp-lane</em><sub>0</sub> only]</b> Warp-wide aggregate reduction of input items (exclusive of the <code>warp_prefix_op</code> value). </td></tr>
    <tr><td class="paramdir">[in,out]</td><td class="paramname">warp_prefix_op</td><td><b>[<em>warp-lane</em><sub>0</sub> only]</b> Call-back functor for specifying a warp-wide prefix to be applied to all inputs. </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="warp__scan_8cuh_source.html#l00679">679</a> of file <a class="el" href="warp__scan_8cuh_source.html">warp_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="aafd14f3920c493b104a3a6a0cc6a39d2"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int LOGICAL_WARPS = 1, int LOGICAL_WARP_THREADS = PtxArchProps::WARP_THREADS&gt; </div>
<div class="memtemplate">
template&lt;typename ScanOp &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_warp_scan.html">cub::WarpScan</a>&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt;::InclusiveScan </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>output</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">ScanOp&#160;</td>
          <td class="paramname"><em>scan_op</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an inclusive prefix sum using the specified binary scan functor in each logical warp. </p>
<p>Supports non-commutative scan operators.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates four concurrent warp-wide inclusive prefix max scans within a block of 128 threads (one per each of the 32-thread warps). </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize WarpScan for 4 warps on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_warp_scan.html" title="The WarpScan class provides collective methods for computing a parallel prefix scan of items partitio...">cub::WarpScan&lt;int, 4&gt;</a> <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for WarpScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> WarpScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain one input item per thread</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data = ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Compute inclusive warp-wide prefix max scans</span></div>
<div class="line">    <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>(temp_storage).InclusiveScan(thread_data, thread_data, <a class="code" href="structcub_1_1_max.html" title="Default max functor. ">cub::Max</a>());</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>0, -1, 2, -3, ..., 126, -127</code>. The corresponding output <code>thread_data</code> in the first warp would be <code>0, 0, 2, 2, ..., 30, 30</code>, the output for the second warp would be <code>32, 32, 34, 34, ..., 62, 62</code>, etc.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ScanOp</td><td><b>[inferred]</b> Binary scan operator type having member <code>T operator()(const T &amp;a, const T &amp;b)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input item. </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output item. May be aliased with <code>input</code>. </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">scan_op</td><td>Binary scan operator </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="warp__scan_8cuh_source.html#l00731">731</a> of file <a class="el" href="warp__scan_8cuh_source.html">warp_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a6a2ac5f642393ccfc698aa515e192a22"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int LOGICAL_WARPS = 1, int LOGICAL_WARP_THREADS = PtxArchProps::WARP_THREADS&gt; </div>
<div class="memtemplate">
template&lt;typename ScanOp &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_warp_scan.html">cub::WarpScan</a>&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt;::InclusiveScan </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>output</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">ScanOp&#160;</td>
          <td class="paramname"><em>scan_op</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>warp_aggregate</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an inclusive prefix sum using the specified binary scan functor in each logical warp. Also provides every thread with the warp-wide <code>warp_aggregate</code> of all inputs. </p>
<p>Supports non-commutative scan operators.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates four concurrent warp-wide inclusive prefix max scans within a block of 128 threads (one per each of the 32-thread warps). </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize WarpScan for 4 warps on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_warp_scan.html" title="The WarpScan class provides collective methods for computing a parallel prefix scan of items partitio...">cub::WarpScan&lt;int, 4&gt;</a> <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for WarpScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> WarpScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain one input item per thread</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data = ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Compute inclusive warp-wide prefix max scans</span></div>
<div class="line">    <span class="keywordtype">int</span> warp_aggregate;</div>
<div class="line">    <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>(temp_storage).InclusiveScan(</div>
<div class="line">        thread_data, thread_data, <a class="code" href="structcub_1_1_max.html" title="Default max functor. ">cub::Max</a>(), warp_aggregate);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>0, -1, 2, -3, ..., 126, -127</code>. The corresponding output <code>thread_data</code> in the first warp would be <code>0, 0, 2, 2, ..., 30, 30</code>, the output for the second warp would be <code>32, 32, 34, 34, ..., 62, 62</code>, etc. Furthermore, <code>warp_aggregate</code> would be assigned <code>30</code> for threads in the first warp, <code>62</code> for threads in the second warp, etc.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ScanOp</td><td><b>[inferred]</b> Binary scan operator type having member <code>T operator()(const T &amp;a, const T &amp;b)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input item. </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output item. May be aliased with <code>input</code>. </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">scan_op</td><td>Binary scan operator </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">warp_aggregate</td><td>Warp-wide aggregate reduction of input items. </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="warp__scan_8cuh_source.html#l00780">780</a> of file <a class="el" href="warp__scan_8cuh_source.html">warp_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="adf74795dbe02947facd24657bc519c40"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int LOGICAL_WARPS = 1, int LOGICAL_WARP_THREADS = PtxArchProps::WARP_THREADS&gt; </div>
<div class="memtemplate">
template&lt;typename ScanOp , typename WarpPrefixOp &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_warp_scan.html">cub::WarpScan</a>&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt;::InclusiveScan </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>output</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">ScanOp&#160;</td>
          <td class="paramname"><em>scan_op</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>warp_aggregate</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">WarpPrefixOp &amp;&#160;</td>
          <td class="paramname"><em>warp_prefix_op</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an inclusive prefix sum using the specified binary scan functor in each logical warp. The call-back functor <code>warp_prefix_op</code> is invoked to provide the "seed" value that logically prefixes the warp's scan inputs. Also provides every thread with the warp-wide <code>warp_aggregate</code> of all inputs. </p>
<p>The <code>warp_prefix_op</code> functor must implement a member function <code>T operator()(T warp_aggregate)</code>. The functor's input parameter <code>warp_aggregate</code> is the same value also returned by the scan operation. The functor will be invoked by the entire warp of threads, however only the return value from <em>lane</em><sub>0</sub> is applied as the threadblock-wide prefix. Can be stateful.</p>
<p>Supports non-commutative scan operators.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates a single thread block of 32 threads (one warp) that progressively computes an inclusive prefix max scan over multiple "tiles" of input using a prefix functor to maintain a running total between block-wide scans. Each tile consists of 32 integer items that are partitioned across the warp. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line"><span class="comment">// A stateful callback functor that maintains a running prefix to be applied</span></div>
<div class="line"><span class="comment">// during consecutive scan operations.</span></div>
<div class="line"><span class="keyword">struct </span>WarpPrefixOp</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Running prefix</span></div>
<div class="line">    <span class="keywordtype">int</span> running_total;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Constructor</span></div>
<div class="line">    __device__ WarpPrefixOp(<span class="keywordtype">int</span> running_total) : running_total(running_total) {}</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Callback operator to be entered by the entire warp. Lane-0 is responsible</span></div>
<div class="line">    <span class="comment">// for returning a value for seeding the warp-wide scan.</span></div>
<div class="line">    __device__ <span class="keywordtype">int</span> operator()(<span class="keywordtype">int</span> warp_aggregate)</div>
<div class="line">    {</div>
<div class="line">        <span class="keywordtype">int</span> old_prefix = running_total;</div>
<div class="line">        running_total = (warp_aggregate &gt; old_prefix) ? warp_aggregate : old_prefix;</div>
<div class="line">        <span class="keywordflow">return</span> old_prefix;</div>
<div class="line">    }</div>
<div class="line">};</div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(<span class="keywordtype">int</span> *d_data, <span class="keywordtype">int</span> num_items, ...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize WarpScan for one warp</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_warp_scan.html" title="The WarpScan class provides collective methods for computing a parallel prefix scan of items partitio...">cub::WarpScan&lt;int, 1&gt;</a> <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for WarpScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> WarpScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Initialize running total</span></div>
<div class="line">    WarpPrefixOp prefix_op(0);</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Have the warp iterate over segments of items</span></div>
<div class="line">    <span class="keywordflow">for</span> (<span class="keywordtype">int</span> block_offset = 0; block_offset &lt; num_items; block_offset += 32)</div>
<div class="line">    {</div>
<div class="line">        <span class="comment">// Load a segment of consecutive items</span></div>
<div class="line">        <span class="keywordtype">int</span> thread_data = d_data[block_offset];</div>
<div class="line"></div>
<div class="line">        <span class="comment">// Collectively compute the warp-wide inclusive prefix max scan</span></div>
<div class="line">        <span class="keywordtype">int</span> warp_aggregate;</div>
<div class="line">        <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>(temp_storage).InclusiveScan(</div>
<div class="line">            thread_data, thread_data, <a class="code" href="structcub_1_1_max.html" title="Default max functor. ">cub::Max</a>(), warp_aggregate, prefix_op);</div>
<div class="line"></div>
<div class="line">        <span class="comment">// Store scanned items to output segment</span></div>
<div class="line">        d_data[block_offset] = thread_data;</div>
<div class="line">    }</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the input <code>d_data</code> is <code>0, -1, 2, -3, 4, -5, ...</code>. The corresponding output for the first segment will be <code>0, 0, 2, 2, ..., 30, 30</code>. The output for the second segment will be <code>32, 32, 34, 34, ..., 62, 62</code>. Furthermore, <code>block_aggregate</code> will be assigned <code>30</code> in all threads after the first scan, assigned <code>62</code> after the second scan, etc.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ScanOp</td><td><b>[inferred]</b> Binary scan operator type having member <code>T operator()(const T &amp;a, const T &amp;b)</code> </td></tr>
    <tr><td class="paramname">WarpPrefixOp</td><td><b>[inferred]</b> Call-back functor type having member <code>T operator()(T warp_aggregate)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input item. </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output item. May be aliased with <code>input</code>. </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">scan_op</td><td>Binary scan operator </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">warp_aggregate</td><td><b>[<em>warp-lane</em><sub>0</sub> only]</b> Warp-wide aggregate reduction of input items (exclusive of the <code>warp_prefix_op</code> value). </td></tr>
    <tr><td class="paramdir">[in,out]</td><td class="paramname">warp_prefix_op</td><td><b>[<em>warp-lane</em><sub>0</sub> only]</b> Call-back functor for specifying a warp-wide prefix to be applied to all inputs. </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="warp__scan_8cuh_source.html#l00869">869</a> of file <a class="el" href="warp__scan_8cuh_source.html">warp_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a64d246eb963f9a18d964e17a6f73fc27"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int LOGICAL_WARPS = 1, int LOGICAL_WARP_THREADS = PtxArchProps::WARP_THREADS&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_warp_scan.html">cub::WarpScan</a>&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt;::InclusiveSum </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>output</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an inclusive prefix sum in each logical warp. </p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates four concurrent warp-wide inclusive prefix sums within a block of 128 threads (one per each of the 32-thread warps). </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize WarpScan for 4 warps on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_warp_scan.html" title="The WarpScan class provides collective methods for computing a parallel prefix scan of items partitio...">cub::WarpScan&lt;int, 4&gt;</a> <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for WarpScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> WarpScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain one input item per thread</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data = ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Compute inclusive warp-wide prefix sums</span></div>
<div class="line">    <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>(temp_storage).InclusiveSum(thread_data, thread_data);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>1, 1, 1, 1, ...</code>. The corresponding output <code>thread_data</code> in each of the four warps of threads will be <code>1, 2, 3, ..., 32</code>. </dd></dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input item. </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output item. May be aliased with <code>input</code>. </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="warp__scan_8cuh_source.html#l00300">300</a> of file <a class="el" href="warp__scan_8cuh_source.html">warp_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a601da756b6729894ec8a52080d920cd4"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int LOGICAL_WARPS = 1, int LOGICAL_WARP_THREADS = PtxArchProps::WARP_THREADS&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_warp_scan.html">cub::WarpScan</a>&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt;::InclusiveSum </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>output</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>warp_aggregate</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an inclusive prefix sum in each logical warp. Also provides every thread with the warp-wide <code>warp_aggregate</code> of all inputs. </p>
<p>The <code>warp_aggregate</code> is undefined in threads other than <em>warp-lane</em><sub>0</sub>.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates four concurrent warp-wide inclusive prefix sums within a block of 128 threads (one per each of the 32-thread warps). </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize WarpScan for 4 warps on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_warp_scan.html" title="The WarpScan class provides collective methods for computing a parallel prefix scan of items partitio...">cub::WarpScan&lt;int, 4&gt;</a> <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for WarpScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> WarpScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain one input item per thread</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data = ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Compute inclusive warp-wide prefix sums</span></div>
<div class="line">    <span class="keywordtype">int</span> warp_aggregate;</div>
<div class="line">    <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>(temp_storage).InclusiveSum(thread_data, thread_data, warp_aggregate);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>1, 1, 1, 1, ...</code>. The corresponding output <code>thread_data</code> in each of the four warps of threads will be <code>1, 2, 3, ..., 32</code>. Furthermore, <code>warp_aggregate</code> for all threads in all warps will be <code>32</code>. </dd></dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input item. </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output item. May be aliased with <code>input</code>. </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">warp_aggregate</td><td>Warp-wide aggregate reduction of input items. </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="warp__scan_8cuh_source.html#l00342">342</a> of file <a class="el" href="warp__scan_8cuh_source.html">warp_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a906d0d6381b3af1ad105ec69547350f3"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int LOGICAL_WARPS = 1, int LOGICAL_WARP_THREADS = PtxArchProps::WARP_THREADS&gt; </div>
<div class="memtemplate">
template&lt;typename WarpPrefixOp &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_warp_scan.html">cub::WarpScan</a>&lt; T, LOGICAL_WARPS, LOGICAL_WARP_THREADS &gt;::InclusiveSum </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>output</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>warp_aggregate</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">WarpPrefixOp &amp;&#160;</td>
          <td class="paramname"><em>warp_prefix_op</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an inclusive prefix sum in each logical warp. Instead of using 0 as the warp-wide prefix, the call-back functor <code>warp_prefix_op</code> is invoked to provide the "seed" value that logically prefixes the warp's scan inputs. Also provides every thread with the warp-wide <code>warp_aggregate</code> of all inputs. </p>
<p>The <code>warp_aggregate</code> is undefined in threads other than <em>warp-lane</em><sub>0</sub>.</p>
<p>The <code>warp_prefix_op</code> functor must implement a member function <code>T operator()(T warp_aggregate)</code>. The functor's input parameter <code>warp_aggregate</code> is the same value also returned by the scan operation. The functor will be invoked by the entire warp of threads, however only the return value from <em>lane</em><sub>0</sub> is applied as the threadblock-wide prefix. Can be stateful.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates a single thread block of 32 threads (one warp) that progressively computes an inclusive prefix sum over multiple "tiles" of input using a prefix functor to maintain a running total between block-wide scans. Each tile consists of 32 integer items that are partitioned across the warp. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line"><span class="comment">// A stateful callback functor that maintains a running prefix to be applied</span></div>
<div class="line"><span class="comment">// during consecutive scan operations.</span></div>
<div class="line"><span class="keyword">struct </span>WarpPrefixOp</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Running prefix</span></div>
<div class="line">    <span class="keywordtype">int</span> running_total;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Constructor</span></div>
<div class="line">    __device__ WarpPrefixOp(<span class="keywordtype">int</span> running_total) : running_total(running_total) {}</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Callback operator to be entered by the entire warp. Lane-0 is responsible</span></div>
<div class="line">    <span class="comment">// for returning a value for seeding the warp-wide scan.</span></div>
<div class="line">    __device__ <span class="keywordtype">int</span> operator()(<span class="keywordtype">int</span> warp_aggregate)</div>
<div class="line">    {</div>
<div class="line">        <span class="keywordtype">int</span> old_prefix = running_total;</div>
<div class="line">        running_total += warp_aggregate;</div>
<div class="line">        <span class="keywordflow">return</span> old_prefix;</div>
<div class="line">    }</div>
<div class="line">};</div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(<span class="keywordtype">int</span> *d_data, <span class="keywordtype">int</span> num_items, ...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize WarpScan for one warp</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_warp_scan.html" title="The WarpScan class provides collective methods for computing a parallel prefix scan of items partitio...">cub::WarpScan&lt;int, 1&gt;</a> <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for WarpScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> WarpScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Initialize running total</span></div>
<div class="line">    WarpPrefixOp prefix_op(0);</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Have the warp iterate over segments of items</span></div>
<div class="line">    <span class="keywordflow">for</span> (<span class="keywordtype">int</span> block_offset = 0; block_offset &lt; num_items; block_offset += 32)</div>
<div class="line">    {</div>
<div class="line">        <span class="comment">// Load a segment of consecutive items</span></div>
<div class="line">        <span class="keywordtype">int</span> thread_data = d_data[block_offset];</div>
<div class="line"></div>
<div class="line">        <span class="comment">// Collectively compute the warp-wide inclusive prefix sum</span></div>
<div class="line">        <span class="keywordtype">int</span> warp_aggregate;</div>
<div class="line">        <a class="code" href="classcub_1_1_warp_scan.html#a27ca414b95f92f2b52701cda45273b05" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">WarpScan</a>(temp_storage).InclusiveSum(</div>
<div class="line">            thread_data, thread_data, warp_aggregate, prefix_op);</div>
<div class="line"></div>
<div class="line">        <span class="comment">// Store scanned items to output segment</span></div>
<div class="line">        d_data[block_offset] = thread_data;</div>
<div class="line">    }</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the input <code>d_data</code> is <code>1, 1, 1, 1, 1, 1, 1, 1, ...</code>. The corresponding output for the first segment will be <code>1, 2, 3, ..., 32</code>. The output for the second segment will be <code>33, 34, 35, ..., 64</code>. Furthermore, the value <code>32</code> will be stored in <code>warp_aggregate</code> for all threads after each scan.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">WarpPrefixOp</td><td><b>[inferred]</b> Call-back functor type having member <code>T operator()(T warp_aggregate)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input item. </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output item. May be aliased with <code>input</code>. </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">warp_aggregate</td><td><b>[<em>warp-lane</em><sub>0</sub> only]</b> Warp-wide aggregate reduction of input items, exclusive of the <code>warp_prefix_op</code> value </td></tr>
    <tr><td class="paramdir">[in,out]</td><td class="paramname">warp_prefix_op</td><td><b>[<em>warp-lane</em><sub>0</sub> only]</b> Call-back functor for specifying a warp-wide prefix to be applied to all inputs. </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="warp__scan_8cuh_source.html#l00426">426</a> of file <a class="el" href="warp__scan_8cuh_source.html">warp_scan.cuh</a>.</p>

</div>
</div>
<hr/>The documentation for this class was generated from the following file:<ul>
<li><a class="el" href="warp__scan_8cuh_source.html">warp_scan.cuh</a></li>
</ul>
</div><!-- contents -->
<!-- HTML footer for doxygen 1.8.3.1-->
<!-- start footer part -->
<hr class="footer"/><address class="footer"><small>
Generated on Fri Aug 23 2013 17:31:13 for CUB by &#160;<a href="http://www.doxygen.org/index.html">
<img class="footer" src="doxygen.png" alt="doxygen"/>
</a> 1.8.4
<br>
&copy; 2013 NVIDIA Corporation
</small></address>
</body>
</html>
